#ifndef _PARTICLE_SYSTEM_CUH_
#define _PARTICLE_SYSTEM_CUH_
/******************************************************************************
    GLOBAL KERNELS
******************************************************************************/
/**
 *  Global Function
 *  Count Pairs Kernel
 *  
 *  -One thread per particle
 *  
 *  This function fills the given array with the number of cells each particle
 * overlaps. These values will be used to fill an array of pairs (Particle, Cell)
 * later on.
 *
 *  @param devPos
 * pointer to the global array containing the positions of each particle.
 *  @param numParticles
 * number of particles currently being simulated.
 *  @param devPairCountArray
 * array of size equals to the number of particles. Will be filled with the
 * number of cells each particle overlaps.
 */
__global__ void kernelCountPairs( float3* devPos, int numParticles, int* devPairCountArray );

/**
 *  Global Function
 *  Fill Pair Arrays
 * 
 *  -One thread per particle
 *
 *  This function fills the "vector of pairs", which is in fact two vectors of
 * integers, one for cell ids and another for particle ids.
 *
 *  @param devPos
 * pointer to the global array containing the positions of each particle.
 *  @param numParticles
 * number of particles currently being simulated.
 *  @param devCountScan
 * array with the offset of the begging of the particle lists of each cell.
 *  @param pairPart
 * particle "side" of the "vector of pairs". Will be filled by this method.
 *  @param pairCell
 * cell "side" of the "vector of pairs". Will be filled by this method.
 */
__global__ void kernelFillPairArrays( float3* devPos, int numParticles, int* devCountScan, int* pairPart, int* pairCell );

/**
 *  Global Function
 *  Fill Cell Limit Array
 *
 *  -One thread per pair (particle, cell)
 *
 *  This function fills the "vector of limits" of the cellls.
 *  The vector limits array is composed by two arrays of integers:
 *  cellLimitBegin and cellLimitEnd.
 *  The first keeps the index in the pair arrays that the cell starts. The later
 *  keeps the index where the next cell starts.
 *
 *  @param pairCell
 * cell "side" of the "vector of pairs". Needs to be ordered by cell id.
 *  @param nPairs
 * number of pairs found.
 *  @param cellLimitBegin
 * keeps the index of the begining of the list of each cell on the pair vector.
 *  @param cellLimitEnd
 * keeps the index of the end of the list of each cell on the pair vector.
 */
__global__ void kernelFillCellLimitArrays( int* pairCell ,int nPairs, int* cellLimitBegin , int* cellLimitEnd );

/**
 *  Global Function
 *  Collide Particles
 *
 *  -One thread per particle
 *
 *  This kernel uses the grid, represented by the arrays particlesOrdered
 * cellLimitBegin and cellLimitEnd, and the current position, velocity and
 * force of the particles, to detect and treat the collisions in system.
 *
 *  @param dt
 * the timestep of the simulation. i.e. how much time passed since last update.
 *  @param devPos
 * pointer to the global array containing the positions of each particle.
 *  @param devVel
 * pointer to the global array containing the velocities of each particle.
 *  @param devForce
 * pointer to the global array containing the forces on each particle.
 *  @param outPos
 * pointer to the global array containing the positions of each particle (output)
 *  @param outVel
 * pointer to the global array containing the velocities of each particle (output)
 *  @param outForce
 * pointer to the global array containing the forces on each particle (output)
 *  @param pairPart
 * particle "side" of the "vector of pairs" ordered by cell id.
 *  @param cellLimitBegin
 * keeps the index of the begining of the list of each cell on the pair vector.
 *  @param cellLimitEnd
 * keeps the index of the end of the list of each cell on the pair vector.
 */
__global__ void kernelCollideParticles( float dt, int numParticles,
                                        float3* devPos, float3* devVel, float3* devForce,
                                        float3* outPos, float3* outVel, float3* outForce,
                                        int* particlesOrdered, int* cellLimitBegin, int* cellLimitEnd );

/**
 *  Global Function
 *  Update Euler
 *
 *  -One thread per paticle
 *
 *  This kernel integrates, by Euler, one step of the simulation.
 *
 *  @param dt
 * the timestep of the simulation. i.e. how much time passed since last update.
 *  @param devPos
 * pointer to the global array containing the positions of each particle.
 *  @param devVel
 * pointer to the global array containing the velocities of each particle.
 *  @param devForce
 * pointer to the global array containing the forces on each particle.
 */
__global__ void kernelUpdateEuler( float dt, float3* devPos, float3* devVel, float3* devForce );

/******************************************************************************
    DEVICE KERNELS
******************************************************************************/
/**
 *  Device Function
 *  Get Collision Vector
 *
 *  Given two particles and it's radius', this funtion returns a vector which
 * has the direction of the colision, from A to B and it has lenght equals to
 * the amount of interpenetrationbetween the two particles.
 *
 *  @param posA
 * position of particle A
 *  @param radiusA
 * radius of particle A
 *  @param posB
 * position of particle B
 *  @param radiusB
 * radius of particle A
 *
 *  @return float3
 * collision vector, or a null vector (0,0,0) if the particles did not touch.
 */
__device__ float3 getCollisionVector( float3 posA, float radiusA, float3 posB, float radiusB );

/**
 *  Device Function
 *  Get Relative Velocity
 *  
 *  Given the velocity of two particles A and B, and the direction from A to B,
 * this function returns the relative velocity between both.
 *
 *  @param velA
 * vectorial velocity of particle A
 *  @param velB
 * vectorial velocity of particle B
 *  @param dir_AB
 * normalize vector in the direction from A to B.
 *
 *  @return float
 * the scalar value of the relative velocity between 2 velocity vectors, given the direction (A to B).
 */
__device__ float getRelativeVelocity( float3 velA, float3 velB, float3 dir_AB );

/**
 *  Device Function
 *  Get Overlaping Cells
 *
 *  This function fills the given array with the ids of all the cells that a
 * particle overlaps.
 *
 *  @param devPos
 * pointer to the global array containing the positions of each particle.
 *  @param pid
 * particle id.
 *  @param cellIds
 * array that will hold the ids of all the cells that the pid particle overlaps
 */
__device__ void getOverlappingCells( float3* devPos, int pid, int* cellIds );

/**
 *  Device Function
 *  Make Cell ID
 *
 *  This functions returns the id of the cell given its integer coordinates.
 *
 *  @param xcell
 * x coordinate of the cell.
 *  @param ycell
 * y coordinate of the cell.
 *  @param zcell
 * z coordinate of the cell.
 *
 *  @return int
 * id of cell.
 */
__device__ int makeCellId( int xcell, int ycell, int zcell );

/**
 *  Device Function
 *  Make Cell ID
 *
 *  This functions returns the id of the cell that contains the given position
 * in world.
 *
 *  @param pos
 * x,y,z position in the world
 *
 *  @return int
 * id of cell.
 */
__device__ int makeCellId( float3 pos );

/**
 *  Device Function
 *  Get Collision Point
 *
 *  This functions returns the 3d position in which particles A and B touched.
 *
 *  @param posA
 * position os the particle A.
 *  @param radiusA
 * radius of particle A.
 *  @param posB
 * position of the particle B.
 *  @param radiusB
 * radius of particle B.
 *
 *  @return float3
 * x,y,z position of the collision.
 */
__device__ float3 getCollisionPoint( float3 posA, float radiusA, float3 posB, float radiusB );

#endif /* _PARTICLE_SYSTEM_CUH_ */
